Skip to content

[HIP] [FlyDSL] Refactor gfx950 A16W16 GEMM with centralized policy selection and tuning - #5145

Merged
coderfeli merged 14 commits into
mainfrom
xyt/flydsl_hgemm_refactor
Sep 3, 2026
Merged

coderfeli merged 14 commits into
mainfrom
xyt/flydsl_hgemm_refactor

Conversation

@xytpai

@xytpai xytpai commented Aug 31, 2026

Copy link
Copy Markdown
Contributor

Summary

  • Rebase gfx950 A16W16 GEMM on the latest FlyDSL example kernel.
  • Replace the legacy splitk_hgemm and small_m_hgemm implementations.
  • Support full-tile/HTI policies, split-K, slice-K, bias, FP32 output, and all four input layouts.
  • Update tuning and occupancy/IOU pruning.
  • Update runtime dispatch, tuned GEMM integration, and AOT compilation.
  • Import the complete upstream gfx950 correctness and benchmark test suite.

Tuned Info

  • Across 2,186 updated model-config entries, the new configurations deliver a 1.055× geometric-mean speedup and a 3.06% median latency reduction.

  • FlyDSL HGEMM is selected for 1,554 of 2,186 entries (71.09%).

Reference

@xytpai
xytpai requested a review from a team August 31, 2026 10:29
@xytpai
xytpai marked this pull request as draft August 31, 2026 10:29
@github-actions

Copy link
Copy Markdown
Contributor

🏷️ CI Guide

Runs automatically on every PR:

  • ✅ Pre-checks (submodule verification, code formatting)
  • ✅ Aiter op tests (gfx942 + gfx950)
  • ✅ Triton tests on MI35X (only when aiter/ops/triton/** or related paths are changed)

Extended tests (opt-in via labels):

Label Tests
ci:gfx1250-ffm-triton Run the five-shard gfx1250 FFM Triton test suite
ci:triton-300x Run an additional Triton test job on MI300X in PRs; main branch always runs both MI35X and MI300X
multigpu Aiter multi-GPU tests on the 8-GPU runner
ci:sglang SGLang integration tests: DeepSeek-R1-MXFP4 accuracy, Qwen 3.5 accuracy
ci:atom ATOM benchmark: DeepSeek-R1-0528, GPT-OSS-120B
ci:atom_full ATOM accuracy suite for PR and main models from ATOM models_accuracy.json
ci:vllm vLLM benchmark: GPT-OSS-120B, DeepSeek-R1-0528, Kimi-K2.5
ci:all All standard extended tests (excludes ci:atom_full)

Only add ci:atom_full for FlyDSL or Triton upgrades.
Add labels via the sidebar or gh pr edit 5145 --add-label <label>

PR title tags & labels:
Component tags ([Triton/Gluon], [HIP], [CK], [ASM], ...) are added to the PR title and as PR labels automatically from the changed files and re-synced on every push — change-type tags like [fix]/[Perf], op tags like [MLA], and human labels (ci:*) are left untouched. Add the no-auto-title label to opt this PR out.

@github-actions github-actions Bot changed the title [FLYDSL] Rebase gfx950 A16W16 GEMM with unified policies [HIP] [FlyDSL] Rebase gfx950 A16W16 GEMM with unified policies Aug 31, 2026
@xytpai
xytpai marked this pull request as ready for review September 2, 2026 17:55
@xytpai xytpai added the ci:atom label Sep 2, 2026
@zufayu
zufayu requested a review from coderfeli September 3, 2026 01:17
@xytpai xytpai changed the title [HIP] [FlyDSL] Rebase gfx950 A16W16 GEMM with unified policies [FLYDSL] Refactor gfx950 A16W16 GEMM with centralized policy selection and tuning Sep 3, 2026
@github-actions github-actions Bot changed the title [FLYDSL] Refactor gfx950 A16W16 GEMM with centralized policy selection and tuning [HIP] [FlyDSL] Refactor gfx950 A16W16 GEMM with centralized policy selection and tuning Sep 3, 2026
@xytpai xytpai removed the HIP label Sep 3, 2026
@github-actions github-actions Bot added the HIP label Sep 3, 2026

@coderfeli coderfeli left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The unit test doesn't follow the style needed, plz clean using unit test skill after this . @xytpai

@coderfeli

coderfeli commented Sep 3, 2026

Copy link
Copy Markdown
Collaborator

Approved and merged with following up requested. Don't forget to fix.

@coderfeli
coderfeli merged commit 4e58a9d into main Sep 3, 2026
58 checks passed
@coderfeli
coderfeli deleted the xyt/flydsl_hgemm_refactor branch September 3, 2026 15:08
jiacao-amd pushed a commit to jiacao-amd/aiter that referenced this pull request Sep 12, 2026
Conflict was kimik3_bf16_tuned_gemm.csv only, and it was textual rather than
semantic. ROCm#5145 retuned 184 rows of that file in place while this branch adds
7; the edits interleave in the same region, so git could not align the hunks.

The two sets do not overlap. Every row ROCm#5145 retuned is outdtype=bfloat16 and
every row this branch adds is outdtype=float32, and outdtype is part of the
index get_GEMM_A16W16_config_ builds, so the two address disjoint keys. The
merged-front GEMM writes an FP32 accumulator it slices router logits out of,
which is why its rows are the fp32 ones.

Resolved by key instead of by hunk: upstream's value for every row upstream
has, this branch's 7 added rows kept, upstream's row order preserved. Verified
that upstream neither reordered nor added/removed rows, that this branch only
appends and never edits a pre-existing row, and that no added key collides with
an upstream key -- if any of those had failed the resolution would not be
mechanical. Result is 536 rows, 536 unique keys, no ragged rows.

ROCm#5145 also refactored tuned_gemm.py, but only along the flydsl branch
(get_flydsl_splitk_hgemm_kernel_params was renamed and the kernel parameter
schema changed). This branch's rows are all libtype=hipblaslt and its code
touches only hipb_create_extension, _hipb_mm and get_GEMM_A16W16_config, whose
signature is unchanged. get_padded_m and the [None, 0, 1] lookup loop are also
untouched, so the M=8 -> padded 16 assumption behind the M=16 row still holds.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants